; ModuleID = 'reduction.cpp'
source_filename = "reduction.cpp"
target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128"
target triple = "x86_64-unknown-linux-gnu"

module asm ".globl _ZSt21ios_base_library_initv"

%"class.std::basic_ostream" = type { ptr, %"class.std::basic_ios" }
%"class.std::basic_ios" = type { %"class.std::ios_base", ptr, i8, i8, ptr, ptr, ptr, ptr }
%"class.std::ios_base" = type { ptr, i64, i64, i32, i32, i32, ptr, %"struct.std::ios_base::_Words", [8 x %"struct.std::ios_base::_Words"], i32, ptr, %"class.std::locale" }
%"struct.std::ios_base::_Words" = type { ptr, i64 }
%"class.std::locale" = type { ptr }
%struct.dim3 = type { i32, i32, i32 }
%struct.hipDeviceProp_tR0600 = type { [256 x i8], %struct.hipUUID_t, [8 x i8], i32, i64, i64, i32, i32, i64, i32, [3 x i32], [3 x i32], i32, i64, i32, i32, i64, i64, i32, i32, i32, i32, i32, i32, i32, i32, i32, [2 x i32], [2 x i32], [3 x i32], [2 x i32], [3 x i32], [3 x i32], i32, [2 x i32], [3 x i32], [2 x i32], i32, [2 x i32], [3 x i32], [2 x i32], [3 x i32], i32, [2 x i32], i64, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i64, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i64, i32, i32, i32, i32, i64, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, [63 x i32], [32 x i32], [256 x i8], i64, i32, %struct.hipDeviceArch_t, ptr, ptr, i32, i32, i32, i32, i32, i32 }
%struct.hipUUID_t = type { [16 x i8] }
%struct.hipDeviceArch_t = type { i24 }
%"class.std::ctype" = type <{ %"class.std::locale::facet.base", [4 x i8], ptr, i8, [7 x i8], ptr, ptr, ptr, i8, [256 x i8], [256 x i8], i8, [6 x i8] }>
%"class.std::locale::facet.base" = type <{ ptr, i32 }>

$__hip_gpubin_handle = comdat any

@_ZSt4cerr = external global %"class.std::basic_ostream", align 8
@.str = private unnamed_addr constant [8 x i8] c"Error: \00", align 1
@_Z23atomic_reduction_kernelPiS_i = constant ptr @_Z38__device_stub__atomic_reduction_kernelPiS_i, align 8
@_Z24atomic_reduction_kernel2PiS_i = constant ptr @_Z39__device_stub__atomic_reduction_kernel2PiS_i, align 8
@_Z24atomic_reduction_kernel3PiS_i = constant ptr @_Z39__device_stub__atomic_reduction_kernel3PiS_i, align 8
@.str.2 = private unnamed_addr constant [25 x i8] c"using default value: %d\0A\00", align 1
@.str.3 = private unnamed_addr constant [15 x i8] c"ARRAYSIZE: %d\0A\00", align 1
@_ZSt4cout = external global %"class.std::basic_ostream", align 8
@.str.4 = private unnamed_addr constant [13 x i8] c"Array size: \00", align 1
@.str.5 = private unnamed_addr constant [4 x i8] c" MB\00", align 1
@.str.6 = private unnamed_addr constant [41 x i8] c"The average performance of reduction is \00", align 1
@.str.7 = private unnamed_addr constant [12 x i8] c" GBytes/sec\00", align 1
@.str.8 = private unnamed_addr constant [32 x i8] c"VERIFICATION: result is CORRECT\00", align 1
@.str.9 = private unnamed_addr constant [36 x i8] c"VERIFICATION: result is INCORRECT!!\00", align 1
@0 = private unnamed_addr constant [33 x i8] c"_Z23atomic_reduction_kernelPiS_i\00", align 1
@1 = private unnamed_addr constant [34 x i8] c"_Z24atomic_reduction_kernel2PiS_i\00", align 1
@2 = private unnamed_addr constant [34 x i8] c"_Z24atomic_reduction_kernel3PiS_i\00", align 1
@__hip_fatbin = external constant i8, section ".hip_fatbin"
@__hip_fatbin_wrapper = internal constant { i32, i32, ptr, ptr } { i32 1212764230, i32 1, ptr @__hip_fatbin, ptr null }, section ".hipFatBinSegment", align 8
@__hip_gpubin_handle = linkonce hidden local_unnamed_addr global ptr null, comdat, align 8
@llvm.global_ctors = appending global [1 x { i32, ptr, ptr }] [{ i32, ptr, ptr } { i32 65535, ptr @__hip_module_ctor, ptr null }]
@str = private unnamed_addr constant [32 x i8] c"Usage: ./reduction num_of_elems\00", align 1

; Function Attrs: mustprogress uwtable
define dso_local void @_Z15check_hip_errorv() local_unnamed_addr #0 {
entry:
  %call = tail call i32 @hipGetLastError()
  %cmp.not = icmp eq i32 %call, 0
  br i1 %cmp.not, label %if.end, label %if.then

if.then:                                          ; preds = %entry
  %call1 = tail call noundef nonnull align 8 dereferenceable(8) ptr @_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc(ptr noundef nonnull align 8 dereferenceable(8) @_ZSt4cerr, ptr noundef nonnull @.str)
  %call2 = tail call ptr @hipGetErrorString(i32 noundef %call)
  %call3 = tail call noundef nonnull align 8 dereferenceable(8) ptr @_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc(ptr noundef nonnull align 8 dereferenceable(8) %call1, ptr noundef %call2)
  %call.i = tail call noundef nonnull align 8 dereferenceable(8) ptr @_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_(ptr noundef nonnull align 8 dereferenceable(8) %call3)
  tail call void @exit(i32 noundef %call) #13
  unreachable

if.end:                                           ; preds = %entry
  ret void
}

; Function Attrs: mustprogress nocallback nofree nosync nounwind willreturn memory(argmem: readwrite)
declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #1

declare i32 @hipGetLastError() local_unnamed_addr #2

; Function Attrs: inlinehint mustprogress uwtable
declare noundef nonnull align 8 dereferenceable(8) ptr @_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc(ptr noundef nonnull align 8 dereferenceable(8), ptr noundef) local_unnamed_addr #3

declare ptr @hipGetErrorString(i32 noundef) local_unnamed_addr #2

; Function Attrs: inlinehint mustprogress uwtable
declare noundef nonnull align 8 dereferenceable(8) ptr @_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_(ptr noundef nonnull align 8 dereferenceable(8)) local_unnamed_addr #3

; Function Attrs: noreturn nounwind
declare void @exit(i32 noundef) local_unnamed_addr #4

; Function Attrs: mustprogress nocallback nofree nosync nounwind willreturn memory(argmem: readwrite)
declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #1

; Function Attrs: norecurse uwtable
define dso_local void @_Z38__device_stub__atomic_reduction_kernelPiS_i(ptr noundef %in, ptr noundef %out, i32 noundef %ARRAYSIZE) #5 {
entry:
  %in.addr = alloca ptr, align 8
  %out.addr = alloca ptr, align 8
  %ARRAYSIZE.addr = alloca i32, align 4
  %grid_dim = alloca %struct.dim3, align 8
  %block_dim = alloca %struct.dim3, align 8
  %shmem_size = alloca i64, align 8
  %stream = alloca ptr, align 8
  store ptr %in, ptr %in.addr, align 8, !tbaa !5
  store ptr %out, ptr %out.addr, align 8, !tbaa !5
  store i32 %ARRAYSIZE, ptr %ARRAYSIZE.addr, align 4, !tbaa !9
  %kernel_args1 = alloca [3 x ptr], align 16
  store ptr %in.addr, ptr %kernel_args1, align 16
  %0 = getelementptr inbounds ptr, ptr %kernel_args1, i64 1
  store ptr %out.addr, ptr %0, align 8
  %1 = getelementptr inbounds ptr, ptr %kernel_args1, i64 2
  store ptr %ARRAYSIZE.addr, ptr %1, align 16
  %2 = call i32 @__hipPopCallConfiguration(ptr nonnull %grid_dim, ptr nonnull %block_dim, ptr nonnull %shmem_size, ptr nonnull %stream)
  %3 = load i64, ptr %shmem_size, align 8
  %4 = load ptr, ptr %stream, align 8
  %grid_dim.coerce.sroa.0.0.copyload = load i64, ptr %grid_dim, align 8
  %grid_dim.coerce.sroa.2.0.grid_dim.sroa_idx = getelementptr inbounds i8, ptr %grid_dim, i64 8
  %grid_dim.coerce.sroa.2.0.copyload = load i32, ptr %grid_dim.coerce.sroa.2.0.grid_dim.sroa_idx, align 8
  %block_dim.coerce.sroa.0.0.copyload = load i64, ptr %block_dim, align 8
  %block_dim.coerce.sroa.2.0.block_dim.sroa_idx = getelementptr inbounds i8, ptr %block_dim, i64 8
  %block_dim.coerce.sroa.2.0.copyload = load i32, ptr %block_dim.coerce.sroa.2.0.block_dim.sroa_idx, align 8
  %call = call noundef i32 @hipLaunchKernel(ptr noundef nonnull @_Z23atomic_reduction_kernelPiS_i, i64 %grid_dim.coerce.sroa.0.0.copyload, i32 %grid_dim.coerce.sroa.2.0.copyload, i64 %block_dim.coerce.sroa.0.0.copyload, i32 %block_dim.coerce.sroa.2.0.copyload, ptr noundef nonnull %kernel_args1, i64 noundef %3, ptr noundef %4)
  ret void
}

declare i32 @__hipPopCallConfiguration(ptr, ptr, ptr, ptr) local_unnamed_addr

declare i32 @hipLaunchKernel(ptr, i64, i32, i64, i32, ptr, i64, ptr) local_unnamed_addr

; Function Attrs: norecurse uwtable
define dso_local void @_Z39__device_stub__atomic_reduction_kernel2PiS_i(ptr noundef %in, ptr noundef %out, i32 noundef %ARRAYSIZE) #5 {
entry:
  %in.addr = alloca ptr, align 8
  %out.addr = alloca ptr, align 8
  %ARRAYSIZE.addr = alloca i32, align 4
  %grid_dim = alloca %struct.dim3, align 8
  %block_dim = alloca %struct.dim3, align 8
  %shmem_size = alloca i64, align 8
  %stream = alloca ptr, align 8
  store ptr %in, ptr %in.addr, align 8, !tbaa !5
  store ptr %out, ptr %out.addr, align 8, !tbaa !5
  store i32 %ARRAYSIZE, ptr %ARRAYSIZE.addr, align 4, !tbaa !9
  %kernel_args1 = alloca [3 x ptr], align 16
  store ptr %in.addr, ptr %kernel_args1, align 16
  %0 = getelementptr inbounds ptr, ptr %kernel_args1, i64 1
  store ptr %out.addr, ptr %0, align 8
  %1 = getelementptr inbounds ptr, ptr %kernel_args1, i64 2
  store ptr %ARRAYSIZE.addr, ptr %1, align 16
  %2 = call i32 @__hipPopCallConfiguration(ptr nonnull %grid_dim, ptr nonnull %block_dim, ptr nonnull %shmem_size, ptr nonnull %stream)
  %3 = load i64, ptr %shmem_size, align 8
  %4 = load ptr, ptr %stream, align 8
  %grid_dim.coerce.sroa.0.0.copyload = load i64, ptr %grid_dim, align 8
  %grid_dim.coerce.sroa.2.0.grid_dim.sroa_idx = getelementptr inbounds i8, ptr %grid_dim, i64 8
  %grid_dim.coerce.sroa.2.0.copyload = load i32, ptr %grid_dim.coerce.sroa.2.0.grid_dim.sroa_idx, align 8
  %block_dim.coerce.sroa.0.0.copyload = load i64, ptr %block_dim, align 8
  %block_dim.coerce.sroa.2.0.block_dim.sroa_idx = getelementptr inbounds i8, ptr %block_dim, i64 8
  %block_dim.coerce.sroa.2.0.copyload = load i32, ptr %block_dim.coerce.sroa.2.0.block_dim.sroa_idx, align 8
  %call = call noundef i32 @hipLaunchKernel(ptr noundef nonnull @_Z24atomic_reduction_kernel2PiS_i, i64 %grid_dim.coerce.sroa.0.0.copyload, i32 %grid_dim.coerce.sroa.2.0.copyload, i64 %block_dim.coerce.sroa.0.0.copyload, i32 %block_dim.coerce.sroa.2.0.copyload, ptr noundef nonnull %kernel_args1, i64 noundef %3, ptr noundef %4)
  ret void
}

; Function Attrs: norecurse uwtable
define dso_local void @_Z39__device_stub__atomic_reduction_kernel3PiS_i(ptr noundef %in, ptr noundef %out, i32 noundef %ARRAYSIZE) #5 {
entry:
  %in.addr = alloca ptr, align 8
  %out.addr = alloca ptr, align 8
  %ARRAYSIZE.addr = alloca i32, align 4
  %grid_dim = alloca %struct.dim3, align 8
  %block_dim = alloca %struct.dim3, align 8
  %shmem_size = alloca i64, align 8
  %stream = alloca ptr, align 8
  store ptr %in, ptr %in.addr, align 8, !tbaa !5
  store ptr %out, ptr %out.addr, align 8, !tbaa !5
  store i32 %ARRAYSIZE, ptr %ARRAYSIZE.addr, align 4, !tbaa !9
  %kernel_args1 = alloca [3 x ptr], align 16
  store ptr %in.addr, ptr %kernel_args1, align 16
  %0 = getelementptr inbounds ptr, ptr %kernel_args1, i64 1
  store ptr %out.addr, ptr %0, align 8
  %1 = getelementptr inbounds ptr, ptr %kernel_args1, i64 2
  store ptr %ARRAYSIZE.addr, ptr %1, align 16
  %2 = call i32 @__hipPopCallConfiguration(ptr nonnull %grid_dim, ptr nonnull %block_dim, ptr nonnull %shmem_size, ptr nonnull %stream)
  %3 = load i64, ptr %shmem_size, align 8
  %4 = load ptr, ptr %stream, align 8
  %grid_dim.coerce.sroa.0.0.copyload = load i64, ptr %grid_dim, align 8
  %grid_dim.coerce.sroa.2.0.grid_dim.sroa_idx = getelementptr inbounds i8, ptr %grid_dim, i64 8
  %grid_dim.coerce.sroa.2.0.copyload = load i32, ptr %grid_dim.coerce.sroa.2.0.grid_dim.sroa_idx, align 8
  %block_dim.coerce.sroa.0.0.copyload = load i64, ptr %block_dim, align 8
  %block_dim.coerce.sroa.2.0.block_dim.sroa_idx = getelementptr inbounds i8, ptr %block_dim, i64 8
  %block_dim.coerce.sroa.2.0.copyload = load i32, ptr %block_dim.coerce.sroa.2.0.block_dim.sroa_idx, align 8
  %call = call noundef i32 @hipLaunchKernel(ptr noundef nonnull @_Z24atomic_reduction_kernel3PiS_i, i64 %grid_dim.coerce.sroa.0.0.copyload, i32 %grid_dim.coerce.sroa.2.0.copyload, i64 %block_dim.coerce.sroa.0.0.copyload, i32 %block_dim.coerce.sroa.2.0.copyload, ptr noundef nonnull %kernel_args1, i64 noundef %3, ptr noundef %4)
  ret void
}

; Function Attrs: norecurse uwtable
define dso_local noundef i32 @main(i32 noundef %argc, ptr nocapture noundef readonly %argv) local_unnamed_addr #5 personality ptr @__gxx_personality_v0 {
entry:
  %in.addr.i = alloca ptr, align 8
  %out.addr.i = alloca ptr, align 8
  %ARRAYSIZE.addr.i = alloca i32, align 4
  %grid_dim.i = alloca %struct.dim3, align 8
  %block_dim.i = alloca %struct.dim3, align 8
  %shmem_size.i = alloca i64, align 8
  %stream.i = alloca ptr, align 8
  %kernel_args1.i = alloca [3 x ptr], align 16
  %in = alloca ptr, align 8
  %out = alloca ptr, align 8
  %props = alloca %struct.hipDeviceProp_tR0600, align 8
  %sum = alloca i32, align 4
  %cmp = icmp slt i32 %argc, 2
  br i1 %cmp, label %if.then, label %if.else

if.then:                                          ; preds = %entry
  %puts = tail call i32 @puts(ptr nonnull dereferenceable(1) @str)
  %call1 = tail call i32 (ptr, ...) @printf(ptr noundef nonnull dereferenceable(1) @.str.2, i32 noundef 52428800)
  br label %if.end

if.else:                                          ; preds = %entry
  %arrayidx = getelementptr inbounds ptr, ptr %argv, i64 1
  %0 = load ptr, ptr %arrayidx, align 8, !tbaa !5
  %call.i = tail call i64 @__isoc23_strtol(ptr noundef nonnull %0, ptr noundef null, i32 noundef 10) #14
  %conv.i = trunc i64 %call.i to i32
  br label %if.end

if.end:                                           ; preds = %if.else, %if.then
  %ARRAYSIZE.0 = phi i32 [ 52428800, %if.then ], [ %conv.i, %if.else ]
  %call3 = tail call i32 (ptr, ...) @printf(ptr noundef nonnull dereferenceable(1) @.str.3, i32 noundef %ARRAYSIZE.0)
  %call1.i = tail call noundef nonnull align 8 dereferenceable(8) ptr @_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l(ptr noundef nonnull align 8 dereferenceable(8) @_ZSt4cout, ptr noundef nonnull @.str.4, i64 noundef 12)
  %conv = zext i32 %ARRAYSIZE.0 to i64
  %mul = shl nuw nsw i64 %conv, 2
  %conv5 = uitofp i64 %mul to double
  %div = fmul contract double %conv5, 0x3F50000000000000
  %div6 = fmul contract double %div, 0x3F50000000000000
  %call.i102 = tail call noundef nonnull align 8 dereferenceable(8) ptr @_ZNSo9_M_insertIdEERSoT_(ptr noundef nonnull align 8 dereferenceable(8) @_ZSt4cout, double noundef %div6)
  %call1.i104 = tail call noundef nonnull align 8 dereferenceable(8) ptr @_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l(ptr noundef nonnull align 8 dereferenceable(8) %call.i102, ptr noundef nonnull @.str.5, i64 noundef 3)
  %vtable.i = load ptr, ptr %call.i102, align 8, !tbaa !11
  %vbase.offset.ptr.i = getelementptr i8, ptr %vtable.i, i64 -24
  %vbase.offset.i = load i64, ptr %vbase.offset.ptr.i, align 8
  %add.ptr.i = getelementptr inbounds i8, ptr %call.i102, i64 %vbase.offset.i
  %_M_ctype.i.i = getelementptr inbounds %"class.std::basic_ios", ptr %add.ptr.i, i64 0, i32 5
  %1 = load ptr, ptr %_M_ctype.i.i, align 8, !tbaa !13
  %tobool.not.i.i.i = icmp eq ptr %1, null
  br i1 %tobool.not.i.i.i, label %if.then.i.i.i, label %_ZSt13__check_facetISt5ctypeIcEERKT_PS3_.exit.i.i

if.then.i.i.i:                                    ; preds = %if.end
  tail call void @_ZSt16__throw_bad_castv() #15
  unreachable

_ZSt13__check_facetISt5ctypeIcEERKT_PS3_.exit.i.i: ; preds = %if.end
  %_M_widen_ok.i.i.i = getelementptr inbounds %"class.std::ctype", ptr %1, i64 0, i32 8
  %2 = load i8, ptr %_M_widen_ok.i.i.i, align 8, !tbaa !22
  %tobool.not.i3.i.i = icmp eq i8 %2, 0
  br i1 %tobool.not.i3.i.i, label %if.end.i.i.i, label %if.then.i4.i.i

if.then.i4.i.i:                                   ; preds = %_ZSt13__check_facetISt5ctypeIcEERKT_PS3_.exit.i.i
  %arrayidx.i.i.i = getelementptr inbounds %"class.std::ctype", ptr %1, i64 0, i32 9, i64 10
  %3 = load i8, ptr %arrayidx.i.i.i, align 1, !tbaa !25
  br label %_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit

if.end.i.i.i:                                     ; preds = %_ZSt13__check_facetISt5ctypeIcEERKT_PS3_.exit.i.i
  tail call void @_ZNKSt5ctypeIcE13_M_widen_initEv(ptr noundef nonnull align 8 dereferenceable(570) %1)
  %vtable.i.i.i = load ptr, ptr %1, align 8, !tbaa !11
  %vfn.i.i.i = getelementptr inbounds ptr, ptr %vtable.i.i.i, i64 6
  %4 = load ptr, ptr %vfn.i.i.i, align 8
  %call.i.i.i = tail call noundef signext i8 %4(ptr noundef nonnull align 8 dereferenceable(570) %1, i8 noundef signext 10)
  br label %_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit

_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit: ; preds = %if.then.i4.i.i, %if.end.i.i.i
  %retval.0.i.i.i = phi i8 [ %3, %if.then.i4.i.i ], [ %call.i.i.i, %if.end.i.i.i ]
  %call1.i168 = tail call noundef nonnull align 8 dereferenceable(8) ptr @_ZNSo3putEc(ptr noundef nonnull align 8 dereferenceable(8) %call.i102, i8 noundef signext %retval.0.i.i.i)
  %call.i.i169 = tail call noundef nonnull align 8 dereferenceable(8) ptr @_ZNSo5flushEv(ptr noundef nonnull align 8 dereferenceable(8) %call1.i168)
  %call12 = tail call noalias ptr @malloc(i64 noundef %mul) #16
  %cmp13273.not = icmp eq i32 %ARRAYSIZE.0, 0
  br i1 %cmp13273.not, label %for.cond.cleanup, label %for.body

for.cond.cleanup:                                 ; preds = %for.body, %_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit
  %checksum.0.lcssa = phi i32 [ 0, %_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit ], [ %add, %for.body ]
  call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %in) #14
  call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %out) #14
  %call.i106 = call i32 @hipMalloc(ptr noundef nonnull %in, i64 noundef %mul)
  %call.i107 = call i32 @hipMalloc(ptr noundef nonnull %out, i64 noundef 4)
  %call.i108 = call i32 @hipGetLastError()
  %cmp.not.i = icmp eq i32 %call.i108, 0
  br i1 %cmp.not.i, label %_Z15check_hip_errorv.exit, label %if.then.i

if.then.i:                                        ; preds = %for.cond.cleanup
  %call1.i109 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc(ptr noundef nonnull align 8 dereferenceable(8) @_ZSt4cerr, ptr noundef nonnull @.str)
  %call2.i = call ptr @hipGetErrorString(i32 noundef %call.i108)
  %call3.i = call noundef nonnull align 8 dereferenceable(8) ptr @_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc(ptr noundef nonnull align 8 dereferenceable(8) %call1.i109, ptr noundef %call2.i)
  %call.i.i110 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_(ptr noundef nonnull align 8 dereferenceable(8) %call3.i)
  call void @exit(i32 noundef %call.i108) #13
  unreachable

_Z15check_hip_errorv.exit:                        ; preds = %for.cond.cleanup
  %5 = load ptr, ptr %in, align 8, !tbaa !5
  %call24 = call i32 @hipMemcpy(ptr noundef %5, ptr noundef %call12, i64 noundef %mul, i32 noundef 1)
  %call25 = call i32 @hipDeviceSynchronize()
  %call.i111 = call i32 @hipGetLastError()
  %cmp.not.i112 = icmp eq i32 %call.i111, 0
  br i1 %cmp.not.i112, label %_Z15check_hip_errorv.exit118, label %if.then.i117

if.then.i117:                                     ; preds = %_Z15check_hip_errorv.exit
  %call1.i113 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc(ptr noundef nonnull align 8 dereferenceable(8) @_ZSt4cerr, ptr noundef nonnull @.str)
  %call2.i114 = call ptr @hipGetErrorString(i32 noundef %call.i111)
  %call3.i115 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc(ptr noundef nonnull align 8 dereferenceable(8) %call1.i113, ptr noundef %call2.i114)
  %call.i.i116 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_(ptr noundef nonnull align 8 dereferenceable(8) %call3.i115)
  call void @exit(i32 noundef %call.i111) #13
  unreachable

_Z15check_hip_errorv.exit118:                     ; preds = %_Z15check_hip_errorv.exit
  call void @llvm.lifetime.start.p0(i64 1472, ptr nonnull %props) #14
  %call26 = call i32 @hipGetDevicePropertiesR0600(ptr noundef nonnull %props, i32 noundef 0)
  %sub = add i32 %ARRAYSIZE.0, 255
  %div28101 = lshr i32 %sub, 8
  %.sroa.speculated = call i32 @llvm.umin.i32(i32 %div28101, i32 2048)
  %call32 = call i64 @_ZNSt6chrono3_V212system_clock3nowEv() #14
  %agg.tmp.sroa.0.0.insert.ext = zext i32 %.sroa.speculated to i64
  %agg.tmp.sroa.0.0.insert.insert = or i64 %agg.tmp.sroa.0.0.insert.ext, 4294967296
  %6 = getelementptr inbounds ptr, ptr %kernel_args1.i, i64 1
  %7 = getelementptr inbounds ptr, ptr %kernel_args1.i, i64 2
  %grid_dim.coerce.sroa.2.0.grid_dim.sroa_idx.i = getelementptr inbounds i8, ptr %grid_dim.i, i64 8
  %block_dim.coerce.sroa.2.0.block_dim.sroa_idx.i = getelementptr inbounds i8, ptr %block_dim.i, i64 8
  br label %for.body38

for.body:                                         ; preds = %_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit, %for.body
  %indvars.iv = phi i64 [ %indvars.iv.next, %for.body ], [ 0, %_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit ]
  %checksum.0275 = phi i32 [ %add, %for.body ], [ 0, %_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit ]
  %call14 = tail call i32 @rand() #14
  %rem = srem i32 %call14, 2
  %arrayidx15 = getelementptr inbounds i32, ptr %call12, i64 %indvars.iv
  store i32 %rem, ptr %arrayidx15, align 4, !tbaa !9
  %add = add nsw i32 %rem, %checksum.0275
  %indvars.iv.next = add nuw nsw i64 %indvars.iv, 1
  %exitcond.not = icmp eq i64 %indvars.iv.next, %conv
  br i1 %exitcond.not, label %for.cond.cleanup, label %for.body, !llvm.loop !26

for.cond35:                                       ; preds = %_Z15check_hip_errorv.exit143
  %inc44 = add nuw nsw i32 %i34.0276, 1
  %exitcond281.not = icmp eq i32 %inc44, 10
  br i1 %exitcond281.not, label %for.cond.cleanup37, label %for.body38, !llvm.loop !28

for.cond.cleanup37:                               ; preds = %for.cond35
  %call47 = call i64 @_ZNSt6chrono3_V212system_clock3nowEv() #14
  %sub.i.i = sub nsw i64 %call47, %call32
  %conv.i.i = sitofp i64 %sub.i.i to double
  %div.i.i = fdiv contract double %conv.i.i, 1.000000e+09
  %conv57 = uitofp i32 %ARRAYSIZE.0 to float
  %mul58 = fmul contract float %conv57, 4.000000e+00
  %mul60 = fmul contract float %mul58, 1.000000e+01
  %call1.i120 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l(ptr noundef nonnull align 8 dereferenceable(8) @_ZSt4cout, ptr noundef nonnull @.str.6, i64 noundef 40)
  %conv62 = fpext float %mul60 to double
  %mul63 = fmul contract double %conv62, 1.000000e-09
  %div64 = fdiv contract double %mul63, %div.i.i
  %call.i121 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZNSo9_M_insertIdEERSoT_(ptr noundef nonnull align 8 dereferenceable(8) @_ZSt4cout, double noundef %div64)
  %call1.i123 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l(ptr noundef nonnull align 8 dereferenceable(8) %call.i121, ptr noundef nonnull @.str.7, i64 noundef 11)
  %vtable.i170 = load ptr, ptr %call.i121, align 8, !tbaa !11
  %vbase.offset.ptr.i171 = getelementptr i8, ptr %vtable.i170, i64 -24
  %vbase.offset.i172 = load i64, ptr %vbase.offset.ptr.i171, align 8
  %add.ptr.i173 = getelementptr inbounds i8, ptr %call.i121, i64 %vbase.offset.i172
  %_M_ctype.i.i174 = getelementptr inbounds %"class.std::basic_ios", ptr %add.ptr.i173, i64 0, i32 5
  %8 = load ptr, ptr %_M_ctype.i.i174, align 8, !tbaa !13
  %tobool.not.i.i.i175 = icmp eq ptr %8, null
  br i1 %tobool.not.i.i.i175, label %if.then.i.i.i176, label %_ZSt13__check_facetISt5ctypeIcEERKT_PS3_.exit.i.i179

if.then.i.i.i176:                                 ; preds = %for.cond.cleanup37
  call void @_ZSt16__throw_bad_castv() #15
  unreachable

_ZSt13__check_facetISt5ctypeIcEERKT_PS3_.exit.i.i179: ; preds = %for.cond.cleanup37
  %_M_widen_ok.i.i.i177 = getelementptr inbounds %"class.std::ctype", ptr %8, i64 0, i32 8
  %9 = load i8, ptr %_M_widen_ok.i.i.i177, align 8, !tbaa !22
  %tobool.not.i3.i.i178 = icmp eq i8 %9, 0
  br i1 %tobool.not.i3.i.i178, label %if.end.i.i.i185, label %if.then.i4.i.i181

if.then.i4.i.i181:                                ; preds = %_ZSt13__check_facetISt5ctypeIcEERKT_PS3_.exit.i.i179
  %arrayidx.i.i.i180 = getelementptr inbounds %"class.std::ctype", ptr %8, i64 0, i32 9, i64 10
  %10 = load i8, ptr %arrayidx.i.i.i180, align 1, !tbaa !25
  br label %_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit189

if.end.i.i.i185:                                  ; preds = %_ZSt13__check_facetISt5ctypeIcEERKT_PS3_.exit.i.i179
  call void @_ZNKSt5ctypeIcE13_M_widen_initEv(ptr noundef nonnull align 8 dereferenceable(570) %8)
  %vtable.i.i.i182 = load ptr, ptr %8, align 8, !tbaa !11
  %vfn.i.i.i183 = getelementptr inbounds ptr, ptr %vtable.i.i.i182, i64 6
  %11 = load ptr, ptr %vfn.i.i.i183, align 8
  %call.i.i.i184 = call noundef signext i8 %11(ptr noundef nonnull align 8 dereferenceable(570) %8, i8 noundef signext 10)
  br label %_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit189

_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit189: ; preds = %if.then.i4.i.i181, %if.end.i.i.i185
  %retval.0.i.i.i186 = phi i8 [ %10, %if.then.i4.i.i181 ], [ %call.i.i.i184, %if.end.i.i.i185 ]
  %call1.i187 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZNSo3putEc(ptr noundef nonnull align 8 dereferenceable(8) %call.i121, i8 noundef signext %retval.0.i.i.i186)
  %call.i.i188 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZNSo5flushEv(ptr noundef nonnull align 8 dereferenceable(8) %call1.i187)
  call void @llvm.lifetime.start.p0(i64 4, ptr nonnull %sum) #14
  %12 = load ptr, ptr %out, align 8, !tbaa !5
  %call68 = call i32 @hipMemcpy(ptr noundef nonnull %sum, ptr noundef %12, i64 noundef 4, i32 noundef 2)
  %call.i125 = call i32 @hipGetLastError()
  %cmp.not.i126 = icmp eq i32 %call.i125, 0
  br i1 %cmp.not.i126, label %_Z15check_hip_errorv.exit132, label %if.then.i131

if.then.i131:                                     ; preds = %_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit189
  %call1.i127 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc(ptr noundef nonnull align 8 dereferenceable(8) @_ZSt4cerr, ptr noundef nonnull @.str)
  %call2.i128 = call ptr @hipGetErrorString(i32 noundef %call.i125)
  %call3.i129 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc(ptr noundef nonnull align 8 dereferenceable(8) %call1.i127, ptr noundef %call2.i128)
  %call.i.i130 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_(ptr noundef nonnull align 8 dereferenceable(8) %call3.i129)
  call void @exit(i32 noundef %call.i125) #13
  unreachable

_Z15check_hip_errorv.exit132:                     ; preds = %_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit189
  %13 = load i32, ptr %sum, align 4, !tbaa !9
  %cmp69 = icmp eq i32 %13, %checksum.0.lcssa
  br i1 %cmp69, label %if.then70, label %if.else74

for.body38:                                       ; preds = %_Z15check_hip_errorv.exit118, %for.cond35
  %i34.0276 = phi i32 [ 0, %_Z15check_hip_errorv.exit118 ], [ %inc44, %for.cond35 ]
  %14 = load ptr, ptr %out, align 8, !tbaa !5
  %call39 = call i32 @hipMemsetAsync(ptr noundef %14, i32 noundef 0, i64 noundef 4, ptr noundef null)
  %call41 = call i32 @__hipPushCallConfiguration(i64 %agg.tmp.sroa.0.0.insert.insert, i32 1, i64 4294967552, i32 1, i64 noundef 0, ptr noundef null)
  %tobool.not = icmp eq i32 %call41, 0
  br i1 %tobool.not, label %kcall.configok, label %do.end

kcall.configok:                                   ; preds = %for.body38
  %15 = load ptr, ptr %in, align 8, !tbaa !5
  %16 = load ptr, ptr %out, align 8, !tbaa !5
  call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %in.addr.i)
  call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %out.addr.i)
  call void @llvm.lifetime.start.p0(i64 4, ptr nonnull %ARRAYSIZE.addr.i)
  call void @llvm.lifetime.start.p0(i64 12, ptr nonnull %grid_dim.i)
  call void @llvm.lifetime.start.p0(i64 12, ptr nonnull %block_dim.i)
  call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %shmem_size.i)
  call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %stream.i)
  call void @llvm.lifetime.start.p0(i64 24, ptr nonnull %kernel_args1.i)
  store ptr %15, ptr %in.addr.i, align 8, !tbaa !5
  store ptr %16, ptr %out.addr.i, align 8, !tbaa !5
  store i32 %ARRAYSIZE.0, ptr %ARRAYSIZE.addr.i, align 4, !tbaa !9
  store ptr %in.addr.i, ptr %kernel_args1.i, align 16
  store ptr %out.addr.i, ptr %6, align 8
  store ptr %ARRAYSIZE.addr.i, ptr %7, align 16
  %17 = call i32 @__hipPopCallConfiguration(ptr nonnull %grid_dim.i, ptr nonnull %block_dim.i, ptr nonnull %shmem_size.i, ptr nonnull %stream.i)
  %18 = load i64, ptr %shmem_size.i, align 8
  %19 = load ptr, ptr %stream.i, align 8
  %grid_dim.coerce.sroa.0.0.copyload.i = load i64, ptr %grid_dim.i, align 8
  %grid_dim.coerce.sroa.2.0.copyload.i = load i32, ptr %grid_dim.coerce.sroa.2.0.grid_dim.sroa_idx.i, align 8
  %block_dim.coerce.sroa.0.0.copyload.i = load i64, ptr %block_dim.i, align 8
  %block_dim.coerce.sroa.2.0.copyload.i = load i32, ptr %block_dim.coerce.sroa.2.0.block_dim.sroa_idx.i, align 8
  %call.i135 = call noundef i32 @hipLaunchKernel(ptr noundef nonnull @_Z23atomic_reduction_kernelPiS_i, i64 %grid_dim.coerce.sroa.0.0.copyload.i, i32 %grid_dim.coerce.sroa.2.0.copyload.i, i64 %block_dim.coerce.sroa.0.0.copyload.i, i32 %block_dim.coerce.sroa.2.0.copyload.i, ptr noundef nonnull %kernel_args1.i, i64 noundef %18, ptr noundef %19)
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %in.addr.i)
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %out.addr.i)
  call void @llvm.lifetime.end.p0(i64 4, ptr nonnull %ARRAYSIZE.addr.i)
  call void @llvm.lifetime.end.p0(i64 12, ptr nonnull %grid_dim.i)
  call void @llvm.lifetime.end.p0(i64 12, ptr nonnull %block_dim.i)
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %shmem_size.i)
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %stream.i)
  call void @llvm.lifetime.end.p0(i64 24, ptr nonnull %kernel_args1.i)
  br label %do.end

do.end:                                           ; preds = %kcall.configok, %for.body38
  %call.i136 = call i32 @hipGetLastError()
  %cmp.not.i137 = icmp eq i32 %call.i136, 0
  br i1 %cmp.not.i137, label %_Z15check_hip_errorv.exit143, label %if.then.i142

if.then.i142:                                     ; preds = %do.end
  %call1.i138 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc(ptr noundef nonnull align 8 dereferenceable(8) @_ZSt4cerr, ptr noundef nonnull @.str)
  %call2.i139 = call ptr @hipGetErrorString(i32 noundef %call.i136)
  %call3.i140 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc(ptr noundef nonnull align 8 dereferenceable(8) %call1.i138, ptr noundef %call2.i139)
  %call.i.i141 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_(ptr noundef nonnull align 8 dereferenceable(8) %call3.i140)
  call void @exit(i32 noundef %call.i136) #13
  unreachable

_Z15check_hip_errorv.exit143:                     ; preds = %do.end
  %call42 = call i32 @hipDeviceSynchronize()
  %call.i144 = call i32 @hipGetLastError()
  %cmp.not.i145 = icmp eq i32 %call.i144, 0
  br i1 %cmp.not.i145, label %for.cond35, label %if.then.i150

if.then.i150:                                     ; preds = %_Z15check_hip_errorv.exit143
  %call1.i146 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc(ptr noundef nonnull align 8 dereferenceable(8) @_ZSt4cerr, ptr noundef nonnull @.str)
  %call2.i147 = call ptr @hipGetErrorString(i32 noundef %call.i144)
  %call3.i148 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc(ptr noundef nonnull align 8 dereferenceable(8) %call1.i146, ptr noundef %call2.i147)
  %call.i.i149 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_(ptr noundef nonnull align 8 dereferenceable(8) %call3.i148)
  call void @exit(i32 noundef %call.i144) #13
  unreachable

if.then70:                                        ; preds = %_Z15check_hip_errorv.exit132
  %call1.i153 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l(ptr noundef nonnull align 8 dereferenceable(8) @_ZSt4cout, ptr noundef nonnull @.str.8, i64 noundef 31)
  %vtable.i190 = load ptr, ptr @_ZSt4cout, align 8, !tbaa !11
  %vbase.offset.ptr.i191 = getelementptr i8, ptr %vtable.i190, i64 -24
  %vbase.offset.i192 = load i64, ptr %vbase.offset.ptr.i191, align 8
  %add.ptr.i193 = getelementptr inbounds i8, ptr @_ZSt4cout, i64 %vbase.offset.i192
  %_M_ctype.i.i194 = getelementptr inbounds %"class.std::basic_ios", ptr %add.ptr.i193, i64 0, i32 5
  %20 = load ptr, ptr %_M_ctype.i.i194, align 8, !tbaa !13
  %tobool.not.i.i.i195 = icmp eq ptr %20, null
  br i1 %tobool.not.i.i.i195, label %if.then.i.i.i196, label %_ZSt13__check_facetISt5ctypeIcEERKT_PS3_.exit.i.i199

if.then.i.i.i196:                                 ; preds = %if.then70
  call void @_ZSt16__throw_bad_castv() #15
  unreachable

_ZSt13__check_facetISt5ctypeIcEERKT_PS3_.exit.i.i199: ; preds = %if.then70
  %_M_widen_ok.i.i.i197 = getelementptr inbounds %"class.std::ctype", ptr %20, i64 0, i32 8
  %21 = load i8, ptr %_M_widen_ok.i.i.i197, align 8, !tbaa !22
  %tobool.not.i3.i.i198 = icmp eq i8 %21, 0
  br i1 %tobool.not.i3.i.i198, label %if.end.i.i.i205, label %if.then.i4.i.i201

if.then.i4.i.i201:                                ; preds = %_ZSt13__check_facetISt5ctypeIcEERKT_PS3_.exit.i.i199
  %arrayidx.i.i.i200 = getelementptr inbounds %"class.std::ctype", ptr %20, i64 0, i32 9, i64 10
  %22 = load i8, ptr %arrayidx.i.i.i200, align 1, !tbaa !25
  br label %_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit209

if.end.i.i.i205:                                  ; preds = %_ZSt13__check_facetISt5ctypeIcEERKT_PS3_.exit.i.i199
  call void @_ZNKSt5ctypeIcE13_M_widen_initEv(ptr noundef nonnull align 8 dereferenceable(570) %20)
  %vtable.i.i.i202 = load ptr, ptr %20, align 8, !tbaa !11
  %vfn.i.i.i203 = getelementptr inbounds ptr, ptr %vtable.i.i.i202, i64 6
  %23 = load ptr, ptr %vfn.i.i.i203, align 8
  %call.i.i.i204 = call noundef signext i8 %23(ptr noundef nonnull align 8 dereferenceable(570) %20, i8 noundef signext 10)
  br label %_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit209

_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit209: ; preds = %if.then.i4.i.i201, %if.end.i.i.i205
  %retval.0.i.i.i206 = phi i8 [ %22, %if.then.i4.i.i201 ], [ %call.i.i.i204, %if.end.i.i.i205 ]
  %call1.i207 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZNSo3putEc(ptr noundef nonnull align 8 dereferenceable(8) @_ZSt4cout, i8 noundef signext %retval.0.i.i.i206)
  %call.i.i208 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZNSo5flushEv(ptr noundef nonnull align 8 dereferenceable(8) %call1.i207)
  %vtable.i210 = load ptr, ptr %call.i.i208, align 8, !tbaa !11
  %vbase.offset.ptr.i211 = getelementptr i8, ptr %vtable.i210, i64 -24
  %vbase.offset.i212 = load i64, ptr %vbase.offset.ptr.i211, align 8
  %add.ptr.i213 = getelementptr inbounds i8, ptr %call.i.i208, i64 %vbase.offset.i212
  %_M_ctype.i.i214 = getelementptr inbounds %"class.std::basic_ios", ptr %add.ptr.i213, i64 0, i32 5
  %24 = load ptr, ptr %_M_ctype.i.i214, align 8, !tbaa !13
  %tobool.not.i.i.i215 = icmp eq ptr %24, null
  br i1 %tobool.not.i.i.i215, label %if.then.i.i.i216, label %_ZSt13__check_facetISt5ctypeIcEERKT_PS3_.exit.i.i219

if.then.i.i.i216:                                 ; preds = %_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit209
  call void @_ZSt16__throw_bad_castv() #15
  unreachable

_ZSt13__check_facetISt5ctypeIcEERKT_PS3_.exit.i.i219: ; preds = %_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit209
  %_M_widen_ok.i.i.i217 = getelementptr inbounds %"class.std::ctype", ptr %24, i64 0, i32 8
  %25 = load i8, ptr %_M_widen_ok.i.i.i217, align 8, !tbaa !22
  %tobool.not.i3.i.i218 = icmp eq i8 %25, 0
  br i1 %tobool.not.i3.i.i218, label %if.end.i.i.i225, label %if.then.i4.i.i221

if.then.i4.i.i221:                                ; preds = %_ZSt13__check_facetISt5ctypeIcEERKT_PS3_.exit.i.i219
  %arrayidx.i.i.i220 = getelementptr inbounds %"class.std::ctype", ptr %24, i64 0, i32 9, i64 10
  %26 = load i8, ptr %arrayidx.i.i.i220, align 1, !tbaa !25
  br label %_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit229

if.end.i.i.i225:                                  ; preds = %_ZSt13__check_facetISt5ctypeIcEERKT_PS3_.exit.i.i219
  call void @_ZNKSt5ctypeIcE13_M_widen_initEv(ptr noundef nonnull align 8 dereferenceable(570) %24)
  %vtable.i.i.i222 = load ptr, ptr %24, align 8, !tbaa !11
  %vfn.i.i.i223 = getelementptr inbounds ptr, ptr %vtable.i.i.i222, i64 6
  %27 = load ptr, ptr %vfn.i.i.i223, align 8
  %call.i.i.i224 = call noundef signext i8 %27(ptr noundef nonnull align 8 dereferenceable(570) %24, i8 noundef signext 10)
  br label %_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit229

_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit229: ; preds = %if.then.i4.i.i221, %if.end.i.i.i225
  %retval.0.i.i.i226 = phi i8 [ %26, %if.then.i4.i.i221 ], [ %call.i.i.i224, %if.end.i.i.i225 ]
  %call1.i227 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZNSo3putEc(ptr noundef nonnull align 8 dereferenceable(8) %call.i.i208, i8 noundef signext %retval.0.i.i.i226)
  br label %if.end78

if.else74:                                        ; preds = %_Z15check_hip_errorv.exit132
  %call1.i157 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l(ptr noundef nonnull align 8 dereferenceable(8) @_ZSt4cout, ptr noundef nonnull @.str.9, i64 noundef 35)
  %vtable.i230 = load ptr, ptr @_ZSt4cout, align 8, !tbaa !11
  %vbase.offset.ptr.i231 = getelementptr i8, ptr %vtable.i230, i64 -24
  %vbase.offset.i232 = load i64, ptr %vbase.offset.ptr.i231, align 8
  %add.ptr.i233 = getelementptr inbounds i8, ptr @_ZSt4cout, i64 %vbase.offset.i232
  %_M_ctype.i.i234 = getelementptr inbounds %"class.std::basic_ios", ptr %add.ptr.i233, i64 0, i32 5
  %28 = load ptr, ptr %_M_ctype.i.i234, align 8, !tbaa !13
  %tobool.not.i.i.i235 = icmp eq ptr %28, null
  br i1 %tobool.not.i.i.i235, label %if.then.i.i.i236, label %_ZSt13__check_facetISt5ctypeIcEERKT_PS3_.exit.i.i239

if.then.i.i.i236:                                 ; preds = %if.else74
  call void @_ZSt16__throw_bad_castv() #15
  unreachable

_ZSt13__check_facetISt5ctypeIcEERKT_PS3_.exit.i.i239: ; preds = %if.else74
  %_M_widen_ok.i.i.i237 = getelementptr inbounds %"class.std::ctype", ptr %28, i64 0, i32 8
  %29 = load i8, ptr %_M_widen_ok.i.i.i237, align 8, !tbaa !22
  %tobool.not.i3.i.i238 = icmp eq i8 %29, 0
  br i1 %tobool.not.i3.i.i238, label %if.end.i.i.i245, label %if.then.i4.i.i241

if.then.i4.i.i241:                                ; preds = %_ZSt13__check_facetISt5ctypeIcEERKT_PS3_.exit.i.i239
  %arrayidx.i.i.i240 = getelementptr inbounds %"class.std::ctype", ptr %28, i64 0, i32 9, i64 10
  %30 = load i8, ptr %arrayidx.i.i.i240, align 1, !tbaa !25
  br label %_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit249

if.end.i.i.i245:                                  ; preds = %_ZSt13__check_facetISt5ctypeIcEERKT_PS3_.exit.i.i239
  call void @_ZNKSt5ctypeIcE13_M_widen_initEv(ptr noundef nonnull align 8 dereferenceable(570) %28)
  %vtable.i.i.i242 = load ptr, ptr %28, align 8, !tbaa !11
  %vfn.i.i.i243 = getelementptr inbounds ptr, ptr %vtable.i.i.i242, i64 6
  %31 = load ptr, ptr %vfn.i.i.i243, align 8
  %call.i.i.i244 = call noundef signext i8 %31(ptr noundef nonnull align 8 dereferenceable(570) %28, i8 noundef signext 10)
  br label %_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit249

_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit249: ; preds = %if.then.i4.i.i241, %if.end.i.i.i245
  %retval.0.i.i.i246 = phi i8 [ %30, %if.then.i4.i.i241 ], [ %call.i.i.i244, %if.end.i.i.i245 ]
  %call1.i247 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZNSo3putEc(ptr noundef nonnull align 8 dereferenceable(8) @_ZSt4cout, i8 noundef signext %retval.0.i.i.i246)
  %call.i.i248 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZNSo5flushEv(ptr noundef nonnull align 8 dereferenceable(8) %call1.i247)
  %vtable.i250 = load ptr, ptr %call.i.i248, align 8, !tbaa !11
  %vbase.offset.ptr.i251 = getelementptr i8, ptr %vtable.i250, i64 -24
  %vbase.offset.i252 = load i64, ptr %vbase.offset.ptr.i251, align 8
  %add.ptr.i253 = getelementptr inbounds i8, ptr %call.i.i248, i64 %vbase.offset.i252
  %_M_ctype.i.i254 = getelementptr inbounds %"class.std::basic_ios", ptr %add.ptr.i253, i64 0, i32 5
  %32 = load ptr, ptr %_M_ctype.i.i254, align 8, !tbaa !13
  %tobool.not.i.i.i255 = icmp eq ptr %32, null
  br i1 %tobool.not.i.i.i255, label %if.then.i.i.i256, label %_ZSt13__check_facetISt5ctypeIcEERKT_PS3_.exit.i.i259

if.then.i.i.i256:                                 ; preds = %_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit249
  call void @_ZSt16__throw_bad_castv() #15
  unreachable

_ZSt13__check_facetISt5ctypeIcEERKT_PS3_.exit.i.i259: ; preds = %_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit249
  %_M_widen_ok.i.i.i257 = getelementptr inbounds %"class.std::ctype", ptr %32, i64 0, i32 8
  %33 = load i8, ptr %_M_widen_ok.i.i.i257, align 8, !tbaa !22
  %tobool.not.i3.i.i258 = icmp eq i8 %33, 0
  br i1 %tobool.not.i3.i.i258, label %if.end.i.i.i265, label %if.then.i4.i.i261

if.then.i4.i.i261:                                ; preds = %_ZSt13__check_facetISt5ctypeIcEERKT_PS3_.exit.i.i259
  %arrayidx.i.i.i260 = getelementptr inbounds %"class.std::ctype", ptr %32, i64 0, i32 9, i64 10
  %34 = load i8, ptr %arrayidx.i.i.i260, align 1, !tbaa !25
  br label %_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit269

if.end.i.i.i265:                                  ; preds = %_ZSt13__check_facetISt5ctypeIcEERKT_PS3_.exit.i.i259
  call void @_ZNKSt5ctypeIcE13_M_widen_initEv(ptr noundef nonnull align 8 dereferenceable(570) %32)
  %vtable.i.i.i262 = load ptr, ptr %32, align 8, !tbaa !11
  %vfn.i.i.i263 = getelementptr inbounds ptr, ptr %vtable.i.i.i262, i64 6
  %35 = load ptr, ptr %vfn.i.i.i263, align 8
  %call.i.i.i264 = call noundef signext i8 %35(ptr noundef nonnull align 8 dereferenceable(570) %32, i8 noundef signext 10)
  br label %_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit269

_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit269: ; preds = %if.then.i4.i.i261, %if.end.i.i.i265
  %retval.0.i.i.i266 = phi i8 [ %34, %if.then.i4.i.i261 ], [ %call.i.i.i264, %if.end.i.i.i265 ]
  %call1.i267 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZNSo3putEc(ptr noundef nonnull align 8 dereferenceable(8) %call.i.i248, i8 noundef signext %retval.0.i.i.i266)
  br label %if.end78

if.end78:                                         ; preds = %_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit269, %_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit229
  %call1.i267.sink = phi ptr [ %call1.i267, %_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit269 ], [ %call1.i227, %_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_.exit229 ]
  %call.i.i268 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZNSo5flushEv(ptr noundef nonnull align 8 dereferenceable(8) %call1.i267.sink)
  %36 = load ptr, ptr %in, align 8, !tbaa !5
  %call79 = call i32 @hipFree(ptr noundef %36)
  %37 = load ptr, ptr %out, align 8, !tbaa !5
  %call80 = call i32 @hipFree(ptr noundef %37)
  %call.i160 = call i32 @hipGetLastError()
  %cmp.not.i161 = icmp eq i32 %call.i160, 0
  br i1 %cmp.not.i161, label %_Z15check_hip_errorv.exit167, label %if.then.i166

if.then.i166:                                     ; preds = %if.end78
  %call1.i162 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc(ptr noundef nonnull align 8 dereferenceable(8) @_ZSt4cerr, ptr noundef nonnull @.str)
  %call2.i163 = call ptr @hipGetErrorString(i32 noundef %call.i160)
  %call3.i164 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc(ptr noundef nonnull align 8 dereferenceable(8) %call1.i162, ptr noundef %call2.i163)
  %call.i.i165 = call noundef nonnull align 8 dereferenceable(8) ptr @_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_(ptr noundef nonnull align 8 dereferenceable(8) %call3.i164)
  call void @exit(i32 noundef %call.i160) #13
  unreachable

_Z15check_hip_errorv.exit167:                     ; preds = %if.end78
  call void @free(ptr noundef %call12) #14
  call void @llvm.lifetime.end.p0(i64 4, ptr nonnull %sum) #14
  call void @llvm.lifetime.end.p0(i64 1472, ptr nonnull %props) #14
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %out) #14
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %in) #14
  ret i32 0
}

; Function Attrs: nofree nounwind
declare noundef i32 @printf(ptr nocapture noundef readonly, ...) local_unnamed_addr #6

; Function Attrs: mustprogress nofree nounwind willreturn allockind("alloc,uninitialized") allocsize(0) memory(inaccessiblemem: readwrite)
declare noalias noundef ptr @malloc(i64 noundef) local_unnamed_addr #7

; Function Attrs: nounwind
declare i32 @rand() local_unnamed_addr #8

declare i32 @hipMemcpy(ptr noundef, ptr noundef, i64 noundef, i32 noundef) local_unnamed_addr #2

declare i32 @hipDeviceSynchronize() local_unnamed_addr #2

declare i32 @hipGetDevicePropertiesR0600(ptr noundef, i32 noundef) local_unnamed_addr #2

; Function Attrs: nounwind
declare i64 @_ZNSt6chrono3_V212system_clock3nowEv() local_unnamed_addr #8

declare i32 @hipMemsetAsync(ptr noundef, i32 noundef, i64 noundef, ptr noundef) local_unnamed_addr #2

declare i32 @__hipPushCallConfiguration(i64, i32, i64, i32, i64 noundef, ptr noundef) local_unnamed_addr #2

declare i32 @hipFree(ptr noundef) local_unnamed_addr #2

; Function Attrs: mustprogress nounwind willreturn allockind("free") memory(argmem: readwrite, inaccessiblemem: readwrite)
declare void @free(ptr allocptr nocapture noundef) local_unnamed_addr #9

; Function Attrs: nounwind
declare i64 @__isoc23_strtol(ptr noundef, ptr noundef, i32 noundef) local_unnamed_addr #8

declare i32 @__gxx_personality_v0(...)

declare noundef nonnull align 8 dereferenceable(8) ptr @_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l(ptr noundef nonnull align 8 dereferenceable(8), ptr noundef, i64 noundef) local_unnamed_addr #2

declare noundef nonnull align 8 dereferenceable(8) ptr @_ZNSo3putEc(ptr noundef nonnull align 8 dereferenceable(8), i8 noundef signext) local_unnamed_addr #2

declare noundef nonnull align 8 dereferenceable(8) ptr @_ZNSo5flushEv(ptr noundef nonnull align 8 dereferenceable(8)) local_unnamed_addr #2

; Function Attrs: noreturn
declare void @_ZSt16__throw_bad_castv() local_unnamed_addr #10

declare void @_ZNKSt5ctypeIcE13_M_widen_initEv(ptr noundef nonnull align 8 dereferenceable(570)) local_unnamed_addr #2

declare noundef nonnull align 8 dereferenceable(8) ptr @_ZNSo9_M_insertIdEERSoT_(ptr noundef nonnull align 8 dereferenceable(8), double noundef) local_unnamed_addr #2

declare i32 @hipMalloc(ptr noundef, i64 noundef) local_unnamed_addr #2

declare i32 @__hipRegisterFunction(ptr, ptr, ptr, ptr, i32, ptr, ptr, ptr, ptr, ptr) local_unnamed_addr

declare ptr @__hipRegisterFatBinary(ptr) local_unnamed_addr

define internal void @__hip_module_ctor() {
entry:
  %0 = load ptr, ptr @__hip_gpubin_handle, align 8
  %1 = icmp eq ptr %0, null
  br i1 %1, label %if, label %exit

if:                                               ; preds = %entry
  %2 = tail call ptr @__hipRegisterFatBinary(ptr nonnull @__hip_fatbin_wrapper)
  store ptr %2, ptr @__hip_gpubin_handle, align 8
  br label %exit

exit:                                             ; preds = %if, %entry
  %3 = phi ptr [ %2, %if ], [ %0, %entry ]
  %4 = tail call i32 @__hipRegisterFunction(ptr %3, ptr nonnull @_Z23atomic_reduction_kernelPiS_i, ptr nonnull @0, ptr nonnull @0, i32 -1, ptr null, ptr null, ptr null, ptr null, ptr null)
  %5 = tail call i32 @__hipRegisterFunction(ptr %3, ptr nonnull @_Z24atomic_reduction_kernel2PiS_i, ptr nonnull @1, ptr nonnull @1, i32 -1, ptr null, ptr null, ptr null, ptr null, ptr null)
  %6 = tail call i32 @__hipRegisterFunction(ptr %3, ptr nonnull @_Z24atomic_reduction_kernel3PiS_i, ptr nonnull @2, ptr nonnull @2, i32 -1, ptr null, ptr null, ptr null, ptr null, ptr null)
  %7 = tail call i32 @atexit(ptr nonnull @__hip_module_dtor)
  ret void
}

declare void @__hipUnregisterFatBinary(ptr) local_unnamed_addr

define internal void @__hip_module_dtor() {
entry:
  %0 = load ptr, ptr @__hip_gpubin_handle, align 8
  %.not = icmp eq ptr %0, null
  br i1 %.not, label %exit, label %if

if:                                               ; preds = %entry
  tail call void @__hipUnregisterFatBinary(ptr nonnull %0)
  store ptr null, ptr @__hip_gpubin_handle, align 8
  br label %exit

exit:                                             ; preds = %if, %entry
  ret void
}

declare i32 @atexit(ptr) local_unnamed_addr

; Function Attrs: nofree nounwind
declare noundef i32 @puts(ptr nocapture noundef readonly) local_unnamed_addr #11

; Function Attrs: nocallback nofree nosync nounwind speculatable willreturn memory(none)
declare i32 @llvm.umin.i32(i32, i32) #12

attributes #0 = { mustprogress uwtable "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #1 = { mustprogress nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) }
attributes #2 = { "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #3 = { inlinehint mustprogress uwtable "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #4 = { noreturn nounwind "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #5 = { norecurse uwtable "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #6 = { nofree nounwind "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #7 = { mustprogress nofree nounwind willreturn allockind("alloc,uninitialized") allocsize(0) memory(inaccessiblemem: readwrite) "alloc-family"="malloc" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #8 = { nounwind "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #9 = { mustprogress nounwind willreturn allockind("free") memory(argmem: readwrite, inaccessiblemem: readwrite) "alloc-family"="malloc" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #10 = { noreturn "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #11 = { nofree nounwind }
attributes #12 = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
attributes #13 = { noreturn nounwind }
attributes #14 = { nounwind }
attributes #15 = { noreturn }
attributes #16 = { nounwind allocsize(0) }

!llvm.module.flags = !{!0, !1, !2, !3}
!llvm.ident = !{!4}

!0 = !{i32 1, !"wchar_size", i32 4}
!1 = !{i32 8, !"PIC Level", i32 2}
!2 = !{i32 7, !"PIE Level", i32 2}
!3 = !{i32 7, !"uwtable", i32 2}
!4 = !{!"clang version 16.0.6 (https://github.com/llvm/llvm-project.git 7cbf1a2591520c2491aa35339f227775f4d3adf6)"}
!5 = !{!6, !6, i64 0}
!6 = !{!"any pointer", !7, i64 0}
!7 = !{!"omnipotent char", !8, i64 0}
!8 = !{!"Simple C++ TBAA"}
!9 = !{!10, !10, i64 0}
!10 = !{!"int", !7, i64 0}
!11 = !{!12, !12, i64 0}
!12 = !{!"vtable pointer", !8, i64 0}
!13 = !{!14, !6, i64 240}
!14 = !{!"_ZTSSt9basic_iosIcSt11char_traitsIcEE", !15, i64 0, !6, i64 216, !7, i64 224, !21, i64 225, !6, i64 232, !6, i64 240, !6, i64 248, !6, i64 256}
!15 = !{!"_ZTSSt8ios_base", !16, i64 8, !16, i64 16, !17, i64 24, !18, i64 28, !18, i64 32, !6, i64 40, !19, i64 48, !7, i64 64, !10, i64 192, !6, i64 200, !20, i64 208}
!16 = !{!"long", !7, i64 0}
!17 = !{!"_ZTSSt13_Ios_Fmtflags", !7, i64 0}
!18 = !{!"_ZTSSt12_Ios_Iostate", !7, i64 0}
!19 = !{!"_ZTSNSt8ios_base6_WordsE", !6, i64 0, !16, i64 8}
!20 = !{!"_ZTSSt6locale", !6, i64 0}
!21 = !{!"bool", !7, i64 0}
!22 = !{!23, !7, i64 56}
!23 = !{!"_ZTSSt5ctypeIcE", !24, i64 0, !6, i64 16, !21, i64 24, !6, i64 32, !6, i64 40, !6, i64 48, !7, i64 56, !7, i64 57, !7, i64 313, !7, i64 569}
!24 = !{!"_ZTSNSt6locale5facetE", !10, i64 8}
!25 = !{!7, !7, i64 0}
!26 = distinct !{!26, !27}
!27 = !{!"llvm.loop.mustprogress"}
!28 = distinct !{!28, !27}
